__attribute__((global, cooperative, reserve_r29_for_scavenge)) void kernel_upsample(int IC, int IH, int IW, int OH, int OW, float *__restrict dev_inp, float *__restrict dev_out)
{
    printf("%d", 1);
    // 获取所有线程数量
    int thread_num = GetThreadNum();
    // 获取当前线程的全局索引
    int thread_idx = GetThreadIdx();

    // 定义与初始化DTE
    tops_dte_ctx_t dte_input[2];
    tops_dte_ctx_t dte_output[2];

    dte_input[0].init();
    dte_input[1].init();
    dte_output[0].init();
    dte_output[1].init();

    tops::event event_input[2];
    tops::event event_output[2]; 
    
    // 选取实际工作线程数
    int c_loops = IC;
    auto nr_working_threads = c_loops <= thread_num ? c_loops : thread_num;

    // 结束多余的工作线程
    if (thread_idx >= nr_working_threads)
    {
        return;
    }

    // 计算每个线程的计算的通道数
    auto base_ic_loops = c_loops / nr_working_threads;
    auto thread_ic_loops = base_ic_loops + ((base_ic_loops * nr_working_threads + thread_idx) < c_loops ? 1 : 0);

    int global_shape[3] = {IC, IH, IW};
    int global_output_shape[3] = {IC, OH, OW};

    int private_shape[3] = {1, IH, IW};
    int private_output_shape[3] = {1, OH, OW};
    auto L1_BUFFER_SIZE = 1 * IH * IW;
    auto L1_BUFFER_SIZE_OUT = 1 * OH * OW;

    // 定义与绑定内存
    tops::mdspan global_input(tops::Global, dev_inp, global_shape);
    tops::mdspan global_output(tops::Global, dev_out, global_output_shape);


    __valigned__ float l1_buffer_input[2][L1_BUFFER_SIZE];
    __valigned__ float l1_buffer_output[2][L1_BUFFER_SIZE_OUT];


    // tops::mdspan l1_input0(tops::Private, l1_buffer_input[0], private_shape);
    // tops::mdspan l1_input1(tops::Private, l1_buffer_input[1], private_shape);
    // tops::mdspan l1_output0(tops::Private, l1_buffer_output[0], private_output_shape);
    // tops::mdspan l1_output1(tops::Private, l1_buffer_output[1], private_output_shape);


    // 全量配置DTE
    for(auto i = 0; i < 2; i++)
    {
        dte_input[i].config_slice(tops::mdspan(tops::Private, l1_buffer_input[i], private_shape), global_input, {thread_idx, 0, 0});
        dte_output[i].config_deslice(global_output, tops::mdspan(tops::Private, l1_buffer_output[i], private_output_shape), {thread_idx, 0, 0});
    }

    event_input[0] = dte_input[0].trigger();
    int iter = 0;

    // 采样过程
    for (auto i = 0; i < thread_ic_loops; i++)
    {
        // ic维度的全局索引
        auto global_ic_index = i * nr_working_threads + thread_idx;
        auto ic_offset_next = (i+1) * nr_working_threads + thread_idx;

        event_input[iter%2].wait();

        if((i+1) < thread_ic_loops)
        {
            dte_input[(iter+1)%2].set_src_offset(0, ic_offset_next);
            event_input[(iter+1)%2] = dte_input[(iter+1)%2].trigger();
        }

        // event_input0 = tops::slice_async(dte_input[0], l1_input0, global_input, {global_ic_index, 0, 0});
        // tops::wait(event_input0);

        // 计算
        for (auto j = 0; j < OH; j++)
        {
            auto global_oh_index = j;
            for (auto k = 0; k < OW; k += 1)
            {
                auto global_ow_index = k;

                // 映射到输出空间
                float ih_f = global_oh_index * (IH - 1) / (float)(OH - 1);
                float iw_f = global_ow_index * (IW - 1) / (float)(OW - 1);

                // 确定输入空间相邻的四个邻近像素
                int ih0 = floorf(ih_f);
                int iw0 = floorf(iw_f);
                int ih1 = ih0 + 1 < IH ? ih0 + 1 : ih0;
                int iw1 = iw0 + 1 < IW ? iw0 + 1 : iw0;

                // 计算小数部分
                float delta_h = ih_f - ih0;
                float delta_w = iw_f - iw0;

                // 计算双线性插值的权重
                float w1 = (1.0f - delta_h) * (1.0f - delta_w);
                float w2 = (1.0f - delta_h) * delta_w;
                float w3 = delta_h * (1.0f - delta_w);
                float w4 = delta_h * delta_w;

                // 计算输入数据的地址
                float val1 = l1_buffer_input[iter%2][ih0 * IW + iw0];
                float val2 = l1_buffer_input[iter%2][ih0 * IW + iw1];
                float val3 = l1_buffer_input[iter%2][ih1 * IW + iw0];
                float val4 = l1_buffer_input[iter%2][ih1 * IW + iw1];

                float out_data = w1 * val1 + w2 * val2 + w3 * val3 + w4 * val4;

                l1_buffer_output[iter%2][global_oh_index * OW + global_ow_index] = out_data;
            }
        }

        if (i != 0)
        {
            event_output[(iter + 1) % 2].wait();
        }

        // 输出
        dte_output[iter%2].set_dst_offset(0, global_ic_index);
        event_output[iter%2] = dte_output[iter%2].trigger();

        if ((i + 1) >= thread_ic_loops)
        {
            event_output[iter % 2].wait();
        }
        iter++;
        // event_output0 = tops::deslice_async(dte_output[0], global_output, l1_output0, {global_ic_index, 0, 0});
        // tops::wait(event_output0);
    }

    // 销毁DTE
    dte_input[0].destroy();
    dte_input[1].destroy();
    dte_output[0].destroy();
    dte_output[1].destroy();
}
void GCU_UpSample(int IC, int IH, int IW, int OH, int OW, float *__restrict dev_inp, float *__restrict dev_out)
{
    static const size_t blocks = 2;
    static const size_t threads = 12;
    // 调用kernel
    kernel_upsample<<<blocks, threads>>>(IC, IH, IW, OH, OW, dev_inp, dev_out);
}